#include <stdio.h>
#include <stdlib.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <fcntl.h>
#include <sys/mman.h>
#include <string.h>
#include <unistd.h>
#include <time.h>
#include <CL/cl.h>
#include <string>
#include <malloc.h>

#define BILLION 1000000000

using namespace std;
extern const char ** createProgram(string, int *);
extern const char * createProgramBinary(string, size_t *);

int main(void){
	int tupleNum = 1024*1024*64;
	int * input = (int *)memalign(256,sizeof(int)*tupleNum);
	int * output = (int *)memalign(256,sizeof(int)*tupleNum);

	for(int i=0;i<tupleNum;i++){
		input[i] = i;
	}
	memset(output,0,sizeof(int)*tupleNum);

	struct timespec start, afBuild, end;

	int psc = 0;
	const char **ps = createProgram("select.cl",&psc);
	//size_t binaryLength = 0;
	//const char * binaryPs = createProgramBinary("a.out", &binaryLength);

	clock_gettime(CLOCK_REALTIME,&start);

	cl_int error = 0;
        cl_device_id device;

        cl_uint numP;
        clGetPlatformIDs(0,NULL,&numP);

        cl_platform_id * pid = new cl_platform_id[numP];
        clGetPlatformIDs(numP, pid, NULL);


        clGetDeviceIDs(pid[0],CL_DEVICE_TYPE_GPU,1,&device,NULL);

        cl_context context = clCreateContext(0,1,&device,NULL,NULL,&error);
        if(error != CL_SUCCESS){
                printf("Error creating context\n");
                exit(-1);
        }

	cl_command_queue_properties prop = 0;
        prop |= CL_QUEUE_PROFILING_ENABLE;

        cl_command_queue queue = clCreateCommandQueue(context, device, prop, &error);
        if(error != CL_SUCCESS){
                printf("Error creating command queue\n");
                exit(-1);
        }

        cl_program program = clCreateProgramWithSource(context, psc, (const char **)ps, 0, 0);
	//int binaryStatus;
        //cl_program tprogram= clCreateProgramWithBinary(context,1,&device,&binaryLength, (const unsigned char **)&binaryPs,&binaryStatus, &error);
        if(error != CL_SUCCESS){
                printf("Error creating program\n");
                exit(-1);
        }

	clBuildProgram(program, 0, 0 , 0 , 0, 0);
        if(error != CL_SUCCESS){
                printf("Error building program\n");
                exit(-1);
        }


	cl_mem column = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int)*tupleNum, NULL, &error);
	cl_mem column1 = clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_ALLOC_HOST_PTR, sizeof(int)*tupleNum, NULL, &error);
	
	void* xx = clEnqueueMapBuffer(queue,column1,CL_TRUE,CL_MAP_WRITE|CL_MAP_READ,0,sizeof(int)*tupleNum,0,0,0,0);
	memcpy(xx,input,sizeof(int)*tupleNum);

	clock_gettime(CLOCK_REALTIME,&afBuild);

	cl_mem result = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int)*tupleNum, NULL, &error);

	cl_event ndrEvt;
	clEnqueueWriteBuffer(queue,column,CL_TRUE,0,sizeof(int)*tupleNum,xx,0,0,&ndrEvt);

	clWaitForEvents(1, &ndrEvt);

	cl_ulong startTime;
	cl_ulong endTime;

	clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_START,
                                         sizeof(cl_ulong),
                                         &startTime,
                                         0);

	clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_END,
                                         sizeof(cl_ulong),
                                         &endTime,
                                         0);

	printf("xxx %f\n",1e-9*(endTime - startTime));

	cl_kernel kernel = clCreateKernel(program, "scan", 0);

	clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&column);
	clSetKernelArg(kernel, 1, sizeof(int), (void *)&tupleNum);
	clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&result);

	size_t gl = 2048*256;
	size_t ll = 256;

	clEnqueueNDRangeKernel(queue, kernel, 1, 0, &gl,&ll,0,0,&ndrEvt);

	clWaitForEvents(1, &ndrEvt);
	clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_START,
                                         sizeof(cl_ulong),
                                         &startTime,
                                         0);

	clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_END,
                                         sizeof(cl_ulong),
                                         &endTime,
                                         0);

	printf("kernel %f\n",1e-9*(endTime - startTime));

	if(error != CL_SUCCESS){
                printf("prefix sum error %d\n",error);
                exit(-1);
        }

	cl_mem output1 = clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_ALLOC_HOST_PTR, sizeof(int)*tupleNum, NULL, &error);
	
	void* yy = clEnqueueMapBuffer(queue,column1,CL_TRUE,CL_MAP_WRITE|CL_MAP_READ,0,sizeof(int)*tupleNum,0,0,0,0);
	
	clEnqueueReadBuffer(queue, result, CL_TRUE, 0, sizeof(int)*(tupleNum), yy,0,0,&ndrEvt);

	clWaitForEvents(1, &ndrEvt);
	clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_START,
                                         sizeof(cl_ulong),
                                         &startTime,
                                         0);

	clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_END,
                                         sizeof(cl_ulong),
                                         &endTime,
                                         0);

	printf("DtoH transfer %f\n",1e-9*(endTime - startTime));



	error = clFinish(queue);
	clReleaseCommandQueue(queue);

	printf("%d\n",output[tupleNum-1]);

	clock_gettime(CLOCK_REALTIME,&end);
        double timeE = (end.tv_sec -  start.tv_sec)* BILLION + end.tv_nsec - start.tv_nsec;
        printf("Table scan CPU time: %lf\n", timeE/1000/1000);

        timeE = (end.tv_sec -  afBuild.tv_sec)* BILLION + end.tv_nsec - afBuild.tv_nsec;
        printf("Table scan CPU time without build: %lf\n", timeE/1000/1000);

}


